/*
    Copyright (C) 2009 Giacomo Spigler

    This program is free software; you can redistribute it and/or modify
    it under the terms of the GNU General Public License as published by
    the Free Software Foundation; either version 2 of the License, or
    (at your option) any later version.

    This program is distributed in the hope that it will be useful,
    but WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
    GNU General Public License for more details.

    You should have received a copy of the GNU General Public License along
    with this program; if not, write to the Free Software Foundation, Inc.,
    51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
*/


#include <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <math.h>

#include <cv.h>

#include <cutil.h>

#include "standard.h"

#include "gabor.h"



texture<unsigned char, 2, cudaReadModeNormalizedFloat> s1_0;
texture<unsigned char, 2, cudaReadModeNormalizedFloat> s1_1;
texture<unsigned char, 2, cudaReadModeNormalizedFloat> s1_2;
texture<unsigned char, 2, cudaReadModeNormalizedFloat> s1_3;



__global__ void C1_GPU(unsigned char *c1_0, unsigned char *c1_1, unsigned char *c1_2, unsigned char *c1_3, int *w, int *h, size_t c1_pitch) {
  unsigned int scale, i, coord;
  int x, y, xi, yi, xi2, yi2, j, k;


//  int w_0=(w[0]-2*GABOR_W-C1_SIZE)/C1_STEP + 1;
  int h_0=(h[0]-2*GABOR_W-C1_SIZE)/C1_STEP + 1;

  for(scale=0; scale<NUMSCALES-1; scale++) {
    int w_out=(w[scale]-2*GABOR_W-C1_SIZE)/C1_STEP + 1;
    int h_out=(h[scale]-2*GABOR_W-C1_SIZE)/C1_STEP + 1;
    int szI=w_out*h_out;

    for(i=blockIdx.x*blockDim.x+threadIdx.x; i<szI; i+=gridDim.x*blockDim.x) {
      //c1 dst coords
      x=i % w_out;
      y=i/w_out;

      xi=x*C1_STEP;
      yi=y*C1_STEP;
      xi2=xi*SCALERATIO;
      yi2=yi*SCALERATIO;
      yi+=scale*h_0;
      yi2+=(scale+1)*h_0;

      float mmax0=0.0;
      float mmax1=0.0;
      float mmax2=0.0;
      float mmax3=0.0;
      //max value over 'scale'(C1_SIZE) and 'scale+1'(C1_SIZE2)
      for(j=0; j<C1_SIZE; j++) {
        for(k=0; k<C1_SIZE; k++) {
          float v0=tex2D(s1_0, xi+j, yi+k);
          float v1=tex2D(s1_1, xi+j, yi+k);
          float v2=tex2D(s1_2, xi+j, yi+k);
          float v3=tex2D(s1_3, xi+j, yi+k);

          mmax0=max(mmax0, v0);
          mmax1=max(mmax1, v1);
          mmax2=max(mmax2, v2);
          mmax3=max(mmax3, v3);

          if(j<C1_SIZE2 && k<C1_SIZE2) {
            v0=tex2D(s1_0, xi2+j, yi2+k);
            v1=tex2D(s1_1, xi2+j, yi2+k);
            v2=tex2D(s1_2, xi2+j, yi2+k);
            v3=tex2D(s1_3, xi2+j, yi2+k);

            mmax0=max(mmax0, v0);
            mmax1=max(mmax1, v1);
            mmax2=max(mmax2, v2);
            mmax3=max(mmax3, v3);
          }

        }
      }


      coord=(y+scale*h_0)*c1_pitch + x;
      c1_0[coord]=mmax0*255.0;
      c1_1[coord]=mmax1*255.0;
      c1_2[coord]=mmax2*255.0;
      c1_3[coord]=mmax3*255.0;


    }

  }

}




extern "C" {



void ComputeC1_2(unsigned char *s1_dev[4], unsigned char *c1_dev[4], int w_host[NUMSCALES], int h_host[NUMSCALES], int *w_dev, int *h_dev, size_t pyr_dev_pitch, size_t c1_pitch,   IplImage *pyr[NUMSCALES]) {
  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>(); //(8, 0, 0, 0, 

  cudaBindTexture2D(0, &s1_0, s1_dev[0], &channelDesc, w_host[0], h_host[0]*NUMSCALES, pyr_dev_pitch);
  cudaBindTexture2D(0, &s1_1, s1_dev[1], &channelDesc, w_host[0], h_host[0]*NUMSCALES, pyr_dev_pitch);
  cudaBindTexture2D(0, &s1_2, s1_dev[2], &channelDesc, w_host[0], h_host[0]*NUMSCALES, pyr_dev_pitch);
  cudaBindTexture2D(0, &s1_3, s1_dev[3], &channelDesc, w_host[0], h_host[0]*NUMSCALES, pyr_dev_pitch);



  C1_GPU<<<C1BLOCKS, C1THREADS>>>(c1_dev[0], c1_dev[1], c1_dev[2], c1_dev[3], w_dev, h_dev, c1_pitch);
  cudaThreadSynchronize();




//TO CHECK!
/*
int ori=0;
int scal=0;

int w_0=(w_host[0]-2*GABOR_W-C1_SIZE)/C1_STEP + 1;
int h_0=(h_host[0]-2*GABOR_W-C1_SIZE)/C1_STEP + 1;
cvZero(pyr[0]);
cudaMemcpy2D(pyr[0]->imageData, pyr[0]->widthStep, c1_dev[ori]+scal*c1_pitch*h_0, c1_pitch, w_0, h_0, cudaMemcpyDeviceToHost);
*/





  cudaUnbindTexture(s1_0);
  cudaUnbindTexture(s1_1);
  cudaUnbindTexture(s1_2);
  cudaUnbindTexture(s1_3);

}








}


